Summary

Fixes a race condition in the GPTQ GEMM kernels that caused incorrect results when input_size > BLOCK_KN_SIZE (128).

Problem

The GPTQ GEMM kernels use multiple thread blocks along the k-dimension (gridDim.z > 1) that accumulate partial results via atomicAdd. The output tensor was being zeroed inside the kernel by blockIdx.z == 0:

if (blockIdx.z == 0) {
  for (int m = 0; m < m_count; m++)
    *((uint64_t*)c_.item_ptr(offset_m + m, n)) = 0;
}
__syncthreads();  // Only syncs within block, not across blocks!

Since __syncthreads() only synchronizes threads within the same block, not across different blocks, this creates a race condition where:

• Blocks with z > 0 may atomicAdd their results before block z=0 finishes zeroing
• Block z=0 may overwrite results that other blocks have already added

This caused numerical errors up to 45x the expected values, particularly when:

• input_size > 128 (triggers multiple k-blocks)
• Errors concentrated in specific rows at m-block boundaries

Solution

1. Pre-zero the output tensor using torch::zeros() instead of torch::empty()
2. Remove the in-kernel zeroing logic from all GPTQ GEMM kernel variants